/*
  Universidade de Sao Paulo - 2014-02
  Programacao Concorrente - Turma 12 Grupo A
  ---
  Bruno Augusto dos Santos - 7696389
  Franco Lamping - 7696351
  Gabriel Massarin de Paula Aguiar - 7573002
*/

#include <iostream>
#include <fstream>
#include <string>
#include <cstdlib>
#include <ctime>

#define INPUT  "in.ppm"
#define OUTPUT "out.ppm"
#define THREADSPERBLOCK_X 16
#define THREADSPERBLOCK_Y 16
#define MASK_SIZE 5

__global__ void applyFilter(unsigned int width,unsigned char *inR,unsigned char *inG,unsigned char *inB, unsigned char *outR,unsigned char *outG,unsigned char *outB) {
  int i = blockIdx.y*blockDim.y+threadIdx.y+2;
  int j = blockIdx.x*blockDim.x+threadIdx.x+2;
  int R = 0, G = 0, B = 0;
  for(int k = -MASK_SIZE/2; k <= MASK_SIZE/2; k++){
    for(int l = -MASK_SIZE/2; l <= MASK_SIZE/2; l++){
      R += inR[(i+k)*width + j + l];
      G += inG[(i+k)*width + j + l];
      B += inB[(i+k)*width + j + l];
    }
  }
  outR[i*width + j] = R/(MASK_SIZE*MASK_SIZE);
  outG[i*width + j] = G/(MASK_SIZE*MASK_SIZE);
  outB[i*width + j] = B/(MASK_SIZE*MASK_SIZE);
}

int main(int argc, char **argv) {
  /* Ponteiros para imagens de entrada e saida no host */
  unsigned char *inR,*inG,*inB, *outR,*outG,*outB;

  /* Ponteiros para imagens de entrada e saida na memoria da GPU */
  unsigned char *dInR, *dInG, *dInB, *dOutR,*dOutG,*dOutB;

  /* Dados sobre a imagem */
  std::string magicNumber;
  unsigned int width, height, depth, pWidth, pHeight;

  std::ifstream inFile;
  std::ofstream outFile;

  inFile.open(INPUT);
  inFile >> magicNumber >> width >> height >> depth;

  /* Verifica se imagem esta no formato permitido */
  if("P3" != magicNumber) {
    std::cout << "Imagem incompativel!" << std::endl;
    return 0;
  }

  pWidth  = width  + MASK_SIZE -1;
  pHeight = height + MASK_SIZE -1;

  /*
  * Calcula o tamanho necessario a adicionar a imagem para que a divisao
  * entre os blocos seja exata
  */
  int paddingW = (THREADSPERBLOCK_X - (pWidth %THREADSPERBLOCK_X))%THREADSPERBLOCK_X;
  int paddingH = (THREADSPERBLOCK_Y - (pHeight%THREADSPERBLOCK_Y))%THREADSPERBLOCK_Y;

  int N = (pHeight+paddingH)*(pWidth+paddingW);

  inR  = new unsigned char[N];
  inG  = new unsigned char[N];
  inB  = new unsigned char[N];
  outR = new unsigned char[N];
  outG = new unsigned char[N];
  outB = new unsigned char[N];

  int r, g, b;
  for(int i = 0; i < pHeight + paddingH; i++){
    for(int j = 0; j < pWidth + paddingW; j++){
      if (i < MASK_SIZE/2 || i >= height+MASK_SIZE/2 || j < MASK_SIZE/2 || j >= width+MASK_SIZE/2) {
        inR[(pWidth+paddingW) * i + j] = 0;
        inG[(pWidth+paddingW) * i + j] = 0;
        inB[(pWidth+paddingW) * i + j] = 0;
      } else {
        inFile >> r >> g >> b;
        inR[(pWidth+paddingW)*i+j] = r;
        inG[(pWidth+paddingW)*i+j] = g;
        inB[(pWidth+paddingW)*i+j] = b;
      }
    }
  }
  inFile.close();

  /* Alocar imagem na memoria da GPU */
  cudaMalloc( (void**)&dInR,  N * sizeof(unsigned char));
  cudaMalloc( (void**)&dInG,  N * sizeof(unsigned char));
  cudaMalloc( (void**)&dInB,  N * sizeof(unsigned char));
  cudaMalloc( (void**)&dOutR, N * sizeof(unsigned char));
  cudaMalloc( (void**)&dOutG, N * sizeof(unsigned char));
  cudaMalloc( (void**)&dOutB, N * sizeof(unsigned char));

  /* Copiar cada componente da imagem para a memoria da GPU */
  cudaMemcpy( dInR, inR, N * sizeof(unsigned char), cudaMemcpyHostToDevice);
  cudaMemcpy( dInG, inG, N * sizeof(unsigned char), cudaMemcpyHostToDevice);
  cudaMemcpy( dInB, inB, N * sizeof(unsigned char), cudaMemcpyHostToDevice);

  dim3 blocksPerGrid  ((pWidth+paddingW)/THREADSPERBLOCK_X, (pHeight+paddingH)/THREADSPERBLOCK_Y,1);
  dim3 threadsPerBlock(THREADSPERBLOCK_X, THREADSPERBLOCK_Y,1);

  /* Registra o inicio do processamento em GPU */
  clock_t t;
  t = clock();

  /* Inicia o processamento */
  applyFilter<<< blocksPerGrid, threadsPerBlock >>>((pWidth+paddingW),dInR,dInG,dInB,dOutR,dOutG,dOutB);

  /* Espera o processamento em GPU finalizar */
  cudaDeviceSynchronize();

  /* Registra o fim do processamento */
  t = clock()-t;
  std::cout << "Elapsed " << t << " clicks (" << ((float)t)/CLOCKS_PER_SEC << " seconds)." << std::endl;

  /* Recupera a imagem da GPU */
  cudaMemcpy(outR, dOutR, N * sizeof(unsigned char), cudaMemcpyDeviceToHost);
  cudaMemcpy(outG, dOutG, N * sizeof(unsigned char), cudaMemcpyDeviceToHost);
  cudaMemcpy(outB, dOutB, N * sizeof(unsigned char), cudaMemcpyDeviceToHost);

  /* Salva a imagem em arquivo, ignorando os paddings de correcao */
  outFile.open(OUTPUT);
  outFile << "P3" << std::endl << width << " " << height << std::endl << depth << std::endl;
  for(int i = 0; i < pHeight; i++){
    for(int j = 0; j < pWidth; j++){
      if (i < MASK_SIZE/2 || i >= height+MASK_SIZE/2 || j < MASK_SIZE/2 || j >= width+MASK_SIZE/2) {
        continue;
      }
      outFile << static_cast<unsigned>(outR[(pWidth+paddingW)*i+j]) << " ";
      outFile << static_cast<unsigned>(outG[(pWidth+paddingW)*i+j]) << " ";
      outFile << static_cast<unsigned>(outB[(pWidth+paddingW)*i+j]) << " ";
    }
    outFile << std::endl;
  }
  outFile.close();

  /* Libera memoria no host e na CPU */
  cudaFree(dInR);
  cudaFree(dInG);
  cudaFree(dInB);
  cudaFree(dOutR);
  cudaFree(dOutG);
  cudaFree(dOutB);
  delete[] inR;
  delete[] inG;
  delete[] inB;
  delete[] outR;
  delete[] outG;
  delete[] outB;

  return 0;
}
